/*
*  This file is part of ygg-brute
*  Copyright (c) 2020 ygg-brute authors
*  See LICENSE for licensing information
*/

#pragma once

__device__ __forceinline__ size_t get_local_id()
{
    return threadIdx.x;
}

__device__ __forceinline__ uint32_t bswap_u32(uint32_t x)
{
    return __byte_perm(x, 0, 0x0123);
}

__device__ __forceinline__ uint64_t bswap_u64(uint64_t x)
{
    return ((uint64_t)bswap_u32(x >> 32)) | (((uint64_t)bswap_u32(x)) << 32);
}

__device__ __forceinline__ size_t global_id()
{
    return blockDim.x * blockIdx.x + threadIdx.x;
}

__device__ __forceinline__ size_t global_size()
{
    return blockDim.x * gridDim.x;
}

__device__ __forceinline__ size_t local_id()
{
    return threadIdx.x;
}

__device__ __forceinline__ size_t local_size()
{
    return blockDim.x;
}

#define SYNCTHREADS() __syncthreads()
#define LOCAL __shared__
#define GLOBAL